// trans.cu

// Copyright (C) 2010  traxex

// Author: 	traxex <traxex@yahoo.cn>
// Created: 	23 Aug 2010
// Version: 	1.0

#include "unittest.h"

const int Y = 256;
const int X = 4096;

matrix<float>	in;
matrix<float>	out;
matrix<float>	out2;


#define BLOCK_DIM 16

__global__ void transpose(float *odata, float *idata, int width, int height)
{
	__shared__ float block[BLOCK_DIM][BLOCK_DIM+1];
	
	// read the matrix tile into shared memory
	unsigned int xIndex = blockIdx.x * BLOCK_DIM + threadIdx.x;
	unsigned int yIndex = blockIdx.y * BLOCK_DIM + threadIdx.y;
	if((xIndex < width) && (yIndex < height))
	{
		unsigned int index_in = yIndex * width + xIndex;
		block[threadIdx.y][threadIdx.x] = idata[index_in];
	}

	__syncthreads();

	// write the transposed matrix tile to global memory
	xIndex = blockIdx.y * BLOCK_DIM + threadIdx.x;
	yIndex = blockIdx.x * BLOCK_DIM + threadIdx.y;
	if((xIndex < height) && (yIndex < width))
	{
		unsigned int index_out = yIndex * height + xIndex;
		odata[index_out] = block[threadIdx.x][threadIdx.y];
	}
}



void performance(){
    LOOPN = 100;
    float (*h)[Y] = new float[X][Y];

    LOOP transpose <<<cusize(Y,X)>>>( out._, in._, X, Y);
    LOOP trans( in, out);
    _timecheck({
            LOOP transpose <<<cusize(Y,X)>>>( out._, in._, X, Y);
        });
    out >> (float*)h;
    _mview( "%6.4f ", h, 20, 20);
    _timecheck({
            LOOP trans( in, out2);
        });
    out2 >> (float*)h;
    _mview( "%6.4f ", h, 20, 20);
}


__global__ void in_init_k( matrix<float> in){
    if( (_my < in.Y) && (_mx < in.X))
        in(_my,_mx) = _my + _mx/1000.0f;
}

bool unittest(){ 
    in  = matrix<float>(Y,X);
    out = matrix<float>(X,Y);
    out2= matrix<float>(X,Y);
    in  = -1;
    in_init_k<<<cusize(Y,X)>>>(in);
    return true;
}
